#include "kernel_operator.h"
#include "kernel_utils.h"
#include "lib/matmul_intf.h"
#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 220
#include "mat_mul_sub_910.h"
#endif
using namespace AscendC;
using namespace matmul;

extern "C" __global__ __aicore__ void mat_mul_sub(GM_ADDR x1, GM_ADDR x2,
                                                        GM_ADDR x3, GM_ADDR y,
                                                        GM_ADDR workspace,
                                                        GM_ADDR tiling) {
  TPipe pipe;
  AscendCUtils::SetOverflow(1);
  KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIC_1_2);
  GM_ADDR user_workspace = GetUserWorkspace(workspace);

  GET_TILING_DATA(tiling_data, tiling);
#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 220
  KernelMatmulSub<DTYPE_X1> op;
#ifdef __DAV_C220_CUBE__
  op.mmObj.SetSubBlockIdx(0);
  op.mmObj.Init(&tiling_data.mmTiling, &pipe);
#endif // __DAV_C220_CUBE__
  op.Init(x1, x2, x3, y, user_workspace, &tiling_data, &pipe);
  op.Process();
#endif // __CCE_AICORE__
}
